




/*
这是关于内存的

共享内存：
是每个SM都有一个的.
Grid
 └── Block (每个 block 最多 1024 个线程)
      └── Warps (每 32 个线程构成一个 warp)
            └── Threads (实际执行指令的最小单元)

GPU
 └── 多个 SM（Streaming Multiprocessor）
      └── 每个 SM 同时调度多个 warp 执行
(
                    ├── SM0 ─┬─ warp0 ─┬─ thread0
                    │        │         ├─ thread1
                    │        │         └─ ...
                    │        └─ warp1 ─ ...
                    ├── SM1 ─ ...

)



在和函数外声明则表示是所有和函数共有的。在内声明则是这个内核里作用
不知道共享空间大小的情况下，可以：
extern __shared__ double s_data[]; 
然后<<<grid, block, share_mem_size（单位：bit）>>>用第三个参数声明（这是cuda独有的语法。）

为了获得高内存带宽，共享内存被分为32个同样大小的内存模型，它们被称为存储体（bank），它们可以被同时访问。有32个存储体是因为在一个线程束中有32个线程。

共享内存是一个一维地址空间

内存存储体的宽度随设备计算能力的不同而变化

字节地址除以4转换为一个4字节字索引（有的储存体宽度是8），然后进行模32操作，将4字节字索引转换为存储体索引。
差不多相当于cache的直接映射了。


当来自相同线程束中的两个线程访问相同的地址时，不会发生存储体冲突。在这种情况下，对于读访问，这个字被广播到请求的线程中；对于写访问，这个字只能由其中一个线程写入，执行这个写入操作的线程是不确定的。


储存体宽度：一次访问时，储体能返回或接收多少字节的数据，也就是说，我们现在讨论的性能问题不是换入换出，而是读取的时候带宽受限
为什么8字节储存体的会比4字节的冲突少？（废话，一次读取8字节数据肯定比一次读取4字节快）
比如字节地址为0、128的两个数据，也就是第0和第32个字，它们是都在bank0里的。8字节的时候可以直接把这两个都读取了。四字节的话，这俩同时请求时，需要排队，（两个总线事务）
只有当两个线程同时访问0（或128）以及256（或512）时，也就是同时访问了bank0的两个不同的8字节数据才会导致冲突。

n向存储体冲突：n个线程访问时发生了上述冲突。解决这种存储体冲突的一个方法是在每N个元素之后添加一个字，这里的N是存储体的数量。
（一个字：一个储存单位。在这里8字节的储存体实际上是装了两个字）


内存填充：通过一定的填充来避免冲突。比如，本来是打算设计一个__shared__ float tile[32][32];的，填充的话可以声明为__shared__ float tile[32][33];



每个SM都有64 KB的片上内存。共享内存和一级缓存共享该硬件资源
cudaDeviceSetCacheConfig()可以设置这俩的占比

参数：
cudaFuncCachePreferNone（默认
cudaFuncCachePreferShared（48kb共享
cudaFuncCachePreferL1（48kbL1
cudaFuncCachePreferEqual（平均

·当核函数使用较多的共享内存时，倾向于更多的共享内存
·当核函数使用更多的寄存器时，倾向于更多的一级缓存





cuda的同步：

在障碍中，所有调用的线程等待其余调用的线程到达障碍点。在内存栅栏中，所有调用的线程必须等到全部内存修改对其余调用线程可见时才能继续执行
为了显式地强制程序以一个确切的顺序执行，必须在应用程序代码中插入内存栅栏和障碍。这是保证与其他线程共享资源的核函数行为正确的唯一途径。(__syncthreads)
这里注意是需要达到相同的障碍点__syncthreads。
如果是(
      if(tid%2 == 0)__syncthreads();
      else __syncthreads();
)
则会死锁


内存栅栏的功能可确保栅栏前的任何内存写操作对栅栏后的其他线程都是可见的。根据所需范围，有3种内存栅栏：块、网格或系统。（
      __threadfence_block
      __threadfence 网格级
      __threadfence_system 系统级
      ）


volatile可以防止编译器优化

全局内存是非常慢的。尽量别直接用……



*/


#include<iostream>
#include<cuda_runtime.h>
#include<vector>
#include "../MyTimer.h"
#include<string>




__global__ void block_crush(int *out, int type)
{
      //extern __shared__ matirx[][]; 这是不合法的
      //如果你在做矩阵运算（如 tile-based GEMM），强烈推荐用一维共享内存数组 + 手动偏移，这样可以通过添加 padding 轻松避免 bank conflict。
      extern __shared__ matirx[]; 
      //现在可以先尝试一下在知道大小可以声明二维的情况下，二维和一维那个快。

      int data;
      if(type == 0)
      {
            //横着来
            data = threadIdx.y * blockDim.x + threadIdx.x;
      }
      else
      {
            data = threadIdx.x * blockDim.y + threadIdx.y;
            //这种情况下，可能正好相隔32个存储体，访问到了同一个存储体，带宽就跟不上了得排队，于是可能会有blocj冲突。
      }
      matirx[data] = data;
      __syncthreads();
      out[data] = matirx[data];

}
void block_crush_run()
{
      //顺便学习一下控制存储体个数的方法吧
      dim3 block(32, 32);
      dim3 grid(1, 1);

      int *out1 = new int[grid.x * grid. y * block.x * block.y];
      int *out2 = new int[grid.x * grid. y * block.x * block.y];

      GPUTimer t;
      t.start();
      block_crush<<<grid, block, sizeof(int) * block.x * block.y>>>(out, 0);
      float ms1 = t.stop();
      
      t.start();
      block_crush<<<grid, block, sizeof(int) * block.x * block.y>>>(out, 1);
      float ms2 = t.stop();


      std::cout<<"row: "<<ms1 << "col: " <<ms2 <<std::endl;


}
int main()
{
      block_crush_run();
}























